#include "md5_computer.h"

__constant__ md5_ctx init_ctx; // Initial MD5 context saved in constant memory.

// __device__ void encode(u8 *output, u32 *input, u32 len) {
// 	u32 i, j;
// 
// 	for(i = 0, j = 0; j < len; i++, j += 4) {
// 		output[j] = (u8) (input[i] & 0xff);
// 		output[j+1] = (u8) ((input[i] >> 8) & 0xff);
// 		output[j+2] = (u8) ((input[i] >> 16) & 0xff);
// 		output[j+3] = (u8) ((input[i] >> 24) & 0xff);
// 	}
// }

// __device__ void md5_memcpy(u8* output, u8* input, u32 len) {
// 	u32 i;
// 
// 	for(i = 0; i < len; i++)
// 		output[i] = input[i];
// }

__device__ size_t custrlen(u8 *str) {
	u8 *s;
	for(s = str; *s != '\0'; ++s);
	return (s - str);
}

__device__ void custrcpy(u8 *dst, u8 *src, size_t size) {
	size_t i;
	
	for(i = 0; i < size && src[i] != '\0'; ++i)
		dst[i] = src[i];
	for(; i < size; ++i)
		dst[i] = '\0';
}

__device__ size_t custrcspn(u8 *str, u8 ch) {
	size_t pos;
	u8 *s;
	u8 sch;

	pos = 0;
	s = str;
	sch = *s;
	while(sch != ch && sch != '\0') {
		++s;
		++pos;
		sch = *s;
	}
	return pos;
}

__device__ void md5_transform(u8 *buffer, size_t len, md5_ctx *ctx) {
	u32 correct_words[16];
	u32 *words = (u32*) buffer;
	size_t nwords = len / sizeof(u32);
	u32 *endp = words + nwords;
	u32 A = ctx->A;
	u32 B = ctx->B;
	u32 C = ctx->C;
	u32 D = ctx->D;

	ctx->total[0] += len;
	if(ctx->total[0] < len)
		++ctx->total[1];
	
	while(words < endp) {
		u32 *cwp = correct_words;
		u32 A_save = A;
		u32 B_save = B;
		u32 C_save = C;
		u32 D_save = D;

#define OP(a, b, c, d, s, T)						\
      do								\
	{								\
	  a += F (b, c, d) + (*cwp++ = *words) + T;		\
	  ++words;							\
	  CYCLIC (a, s);						\
	  a += b;							\
	}								\
      while (0)

      /* It is unfortunate that C does not provide an operator for
	 cyclic rotation.  Hope the C compiler is smart enough.  */
#define CYCLIC(w, s) (w = (w << s) | (w >> (32 - s)))

      /* Before we start, one word to the strange constants.
	 They are defined in RFC 1321 as

	 T[i] = (int) (4294967296.0 * fabs (sin (i))), i=1..64
       */

      /* Round 1.  */
      OP (A, B, C, D,  7, 0xd76aa478);
      OP (D, A, B, C, 12, 0xe8c7b756);
      OP (C, D, A, B, 17, 0x242070db);
      OP (B, C, D, A, 22, 0xc1bdceee);
      OP (A, B, C, D,  7, 0xf57c0faf);
      OP (D, A, B, C, 12, 0x4787c62a);
      OP (C, D, A, B, 17, 0xa8304613);
      OP (B, C, D, A, 22, 0xfd469501);
      OP (A, B, C, D,  7, 0x698098d8);
      OP (D, A, B, C, 12, 0x8b44f7af);
      OP (C, D, A, B, 17, 0xffff5bb1);
      OP (B, C, D, A, 22, 0x895cd7be);
      OP (A, B, C, D,  7, 0x6b901122);
      OP (D, A, B, C, 12, 0xfd987193);
      OP (C, D, A, B, 17, 0xa679438e);
      OP (B, C, D, A, 22, 0x49b40821);

      /* For the second to fourth round we have the possibly swapped words
	 in CORRECT_WORDS.  Redefine the macro to take an additional first
	 argument specifying the function to use.  */
#undef OP
#define OP(f, a, b, c, d, k, s, T)					\
      do 								\
	{								\
	  a += f (b, c, d) + correct_words[k] + T;			\
	  CYCLIC (a, s);						\
	  a += b;							\
	}								\
      while (0)

      /* Round 2.  */
      OP (G, A, B, C, D,  1,  5, 0xf61e2562);
      OP (G, D, A, B, C,  6,  9, 0xc040b340);
      OP (G, C, D, A, B, 11, 14, 0x265e5a51);
      OP (G, B, C, D, A,  0, 20, 0xe9b6c7aa);
      OP (G, A, B, C, D,  5,  5, 0xd62f105d);
      OP (G, D, A, B, C, 10,  9, 0x02441453);
      OP (G, C, D, A, B, 15, 14, 0xd8a1e681);
      OP (G, B, C, D, A,  4, 20, 0xe7d3fbc8);
      OP (G, A, B, C, D,  9,  5, 0x21e1cde6);
      OP (G, D, A, B, C, 14,  9, 0xc33707d6);
      OP (G, C, D, A, B,  3, 14, 0xf4d50d87);
      OP (G, B, C, D, A,  8, 20, 0x455a14ed);
      OP (G, A, B, C, D, 13,  5, 0xa9e3e905);
      OP (G, D, A, B, C,  2,  9, 0xfcefa3f8);
      OP (G, C, D, A, B,  7, 14, 0x676f02d9);
      OP (G, B, C, D, A, 12, 20, 0x8d2a4c8a);

      /* Rund 3.  */
      OP (H, A, B, C, D,  5,  4, 0xfffa3942);
      OP (H, D, A, B, C,  8, 11, 0x8771f681);
      OP (H, C, D, A, B, 11, 16, 0x6d9d6122);
      OP (H, B, C, D, A, 14, 23, 0xfde5380c);
      OP (H, A, B, C, D,  1,  4, 0xa4beea44);
      OP (H, D, A, B, C,  4, 11, 0x4bdecfa9);
      OP (H, C, D, A, B,  7, 16, 0xf6bb4b60);
      OP (H, B, C, D, A, 10, 23, 0xbebfbc70);
      OP (H, A, B, C, D, 13,  4, 0x289b7ec6);
      OP (H, D, A, B, C,  0, 11, 0xeaa127fa);
      OP (H, C, D, A, B,  3, 16, 0xd4ef3085);
      OP (H, B, C, D, A,  6, 23, 0x04881d05);
      OP (H, A, B, C, D,  9,  4, 0xd9d4d039);
      OP (H, D, A, B, C, 12, 11, 0xe6db99e5);
      OP (H, C, D, A, B, 15, 16, 0x1fa27cf8);
      OP (H, B, C, D, A,  2, 23, 0xc4ac5665);

      /* Rund 4.  */
      OP (I, A, B, C, D,  0,  6, 0xf4292244);
      OP (I, D, A, B, C,  7, 10, 0x432aff97);
      OP (I, C, D, A, B, 14, 15, 0xab9423a7);
      OP (I, B, C, D, A,  5, 21, 0xfc93a039);
      OP (I, A, B, C, D, 12,  6, 0x655b59c3);
      OP (I, D, A, B, C,  3, 10, 0x8f0ccc92);
      OP (I, C, D, A, B, 10, 15, 0xffeff47d);
      OP (I, B, C, D, A,  1, 21, 0x85845dd1);
      OP (I, A, B, C, D,  8,  6, 0x6fa87e4f);
      OP (I, D, A, B, C, 15, 10, 0xfe2ce6e0);
      OP (I, C, D, A, B,  6, 15, 0xa3014314);
      OP (I, B, C, D, A, 13, 21, 0x4e0811a1);
      OP (I, A, B, C, D,  4,  6, 0xf7537e82);
      OP (I, D, A, B, C, 11, 10, 0xbd3af235);
      OP (I, C, D, A, B,  2, 15, 0x2ad7d2bb);
      OP (I, B, C, D, A,  9, 21, 0xeb86d391);

      /* Add the starting values of the context.  */
      A += A_save;
      B += B_save;
      C += C_save;
      D += D_save;
    }

  /* Put checksum in context given as argument.  */
  ctx->A = A;
  ctx->B = B;
  ctx->C = C;
  ctx->D = D;
}

// __device__ void md5_init(md5_ctx *ctx) {
// 	ctx->total[0] = ctx->total[1] = 0;
// 	ctx->A = 0x67452301;
// 	ctx->B = 0xefcdab89;
// 	ctx->C = 0x98badcfe;
// 	ctx->D = 0x10325476;
// 	ctx->buflen = 0;
// 
// 	int i;
// 	for(i = 0; i < 128; ++i)
// 		ctx->buffer[i] = 0;
// }

__device__ void md5_update(md5_ctx *ctx, u8 *buffer, u32 len) {
	u32 temp;
	u32 copy_index;
	temp = ctx->buflen;
	if(temp != 0) {
		size_t left_over = temp;
		size_t add;
		if(128 - left_over > len)
			add = len;
		else
			add = 128 - left_over;

		/* That is: md5_memcpy(&ctx->buffer[left_over], buffer, add); */
		for(copy_index = 0; copy_index < add; ++copy_index) {
			ctx->buffer[left_over + copy_index] = buffer[copy_index];
		}
		temp += add;

		if(temp > 64) {
			md5_transform(ctx->buffer, temp & ~63, ctx);

			temp &= 63;
			/* That is: md5_memcpy(ctx->buffer, &ctx->buffer[(left_over + add) & ~63], temp); */
			for(copy_index = 0; copy_index < temp; ++copy_index) {
				ctx->buffer[copy_index] = ctx->buffer[((left_over + add) & ~63) + copy_index];
			}
		}

		buffer = buffer + add;
		len -= add;
	}
	if(len >= 64) {
	#define UNALIGNED_P(p) (((uptr) p) % sizeof(u32) != 0)
		if(UNALIGNED_P(buffer))
			while(len > 64) {
				/* That is: md5_memcpy(ctx->buffer, buffer, 64); */
				for(copy_index = 0; copy_index < 64; ++copy_index) {
					ctx->buffer[copy_index] = buffer[copy_index];
				}
				md5_transform(ctx->buffer, 64, ctx);
				buffer = buffer + 64;
				len -= 64;
			}
		else {
			md5_transform(buffer, len & ~63, ctx);
			buffer += (len & ~63);
			len &= 63;
		}
	}
	if(len > 0) {
		size_t left_over = temp;
		/* That is: md5_memcpy(&ctx->buffer[left_over], buffer, len); */
		for(copy_index = 0; copy_index < len; ++copy_index) {
			ctx->buffer[left_over + copy_index] = buffer[copy_index];
		}
		left_over += len;
		if(left_over >= 64) {
			md5_transform(ctx->buffer, 64, ctx);
			left_over -= 64;
			/* That is: md5_memcpy(ctx->buffer, &ctx->buffer[64], left_over); */
			for(copy_index = 0; copy_index < left_over; ++copy_index) {
				ctx->buffer[copy_index] = ctx->buffer[64 + copy_index];
			}
		}
		temp = left_over;
	}
	ctx->buflen = temp;
}

__device__ void md5_final(u8 digest[16], md5_ctx *ctx) {
	u32 bytes = ctx->buflen;
	u32 state[4];
	u32 copy_index;
	u32 i, j;
	size_t pad;
	u8 PADDING[64] = {
		0x80, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
		0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,
		0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0
	};


	ctx->total[0] += bytes;
	if(ctx->total[0] < bytes)
		++ctx->total[1];
	
	if(bytes >= 56)
		pad = 120 - bytes;
	else
		pad = 56 - bytes;
	
	/* That is: md5_memcpy(&ctx->buffer[bytes], PADDING, pad); */
	for(copy_index = 0; copy_index < pad; ++copy_index) {
		ctx->buffer[bytes + copy_index] = PADDING[copy_index];
	}

	ctx->buffer32[(bytes + pad)/4] = (ctx->total[0] << 3);
	ctx->buffer32[(bytes + pad + 4)/4] = ((ctx->total[1] << 3) | (ctx->total[0] >> 29));

	md5_transform(ctx->buffer, bytes + pad + 8, ctx);

	state[0] = ctx->A;
	state[1] = ctx->B;
	state[2] = ctx->C;
	state[3] = ctx->D;

	/* That is: encode(digest, state, 16); */
	for(i = 0, j = 0; j < 16; i++, j += 4) {
		digest[j] = (u8) (state[i] & 0xff);
		digest[j+1] = (u8) ((state[i] >> 8) & 0xff);
		digest[j+2] = (u8) ((state[i] >> 16) & 0xff);
		digest[j+3] = (u8) ((state[i] >> 24) & 0xff);
	}
}

/* Function that implements the GNU+Linux crypt library function using the MD5 hash algorithm. It computes the hash of input parameter 'string' using salt '__salt' and saves the resulting hash into 'digest'. */
__device__ void compute_md5(u8 *string, u8 *__salt, u8 *digest) {
	u8 key[MAX_STR_LENGTH];
	u8 ____salt[MAX_SALT_LENGTH];
	u8 md5_salt_prefix[] = "$1$";
	u8 *salt;
	size_t key_len;
	size_t salt_len;
	size_t cnt;
	md5_ctx ctx;
	md5_ctx alt_ctx;
	u8 alt_digest[16];

	/* Copy input string into kernel memory. */
	custrcpy(key, string, MAX_STR_LENGTH);
	key_len = custrlen(key);
	
	/* Copy salt string into kernel memory. */
	custrcpy(____salt, __salt, MAX_SALT_LENGTH);
	salt = ____salt;

	/* Get salt length.
		N.B. The salt is considered without '$1$' prefix, since this is an MD5 breaker. */
	salt_len = custrcspn(salt, '$');
	if(salt_len  > 8)
		salt_len = 8;

	/* Prepare for the real work.  */
	ctx = init_ctx;

	/* Add the key string.  */
	md5_update(&ctx, key, key_len);

	// /* Because the SALT argument need not always have the salt prefix we
	//  add it separately.  */
	md5_update(&ctx, md5_salt_prefix, sizeof(md5_salt_prefix) - 1);

	// /* The last part is the salt string.  This must be at most 8
	//  characters and it ends at the first `$' character (for
	//  compatibility with existing implementations).  */
	md5_update(&ctx, salt, salt_len);

	/* Compute alternate MD5 sum with input KEY, SALT, and KEY.  The
	 final result will be added to the first context.  */
	alt_ctx = init_ctx;

	/* Add key.  */
	md5_update(&alt_ctx, key, key_len);


	/* Add salt.  */
	md5_update(&alt_ctx, salt, salt_len);

	/* Add key again.  */
	md5_update(&alt_ctx, key, key_len);

	/* Now get result of this (16 bytes) and add it to the other
	 context.  */
	md5_final(alt_digest, &alt_ctx);

	/* Add for any character in the key one byte of the alternate sum.  */
	for(cnt = key_len; cnt > 16; cnt -= 16)
		md5_update(&ctx, alt_digest, 16);
	md5_update(&ctx, alt_digest, cnt);

	/* For the following code we need a NUL byte.  */
	*alt_digest = '\0';

	/* The original implementation now does something weird: for every 1
	 bit in the key the first 0 is added to the buffer, for every 0
	 bit the first character of the key.  This does not seem to be
	 what was intended but we have to follow this to be compatible.  */
	for(cnt = key_len; cnt > 0; cnt >>= 1) {
		if((cnt & 1) != 0)
			md5_update(&ctx, alt_digest, 1);
		else
			md5_update(&ctx, key, 1);
	}

	/* Create intermediate result.  */
	md5_final(alt_digest, &ctx);

	/* Now comes another weirdness.  In fear of password crackers here
	 comes a quite long loop which just processes the output of the
	 previous round again.  We cannot ignore this here.  */
	for(cnt = 0; cnt < 1000; ++cnt) {
		/* New context.  */
		ctx = init_ctx;

		/* Add key or last result.  */
		if((cnt & 1) != 0)
			md5_update(&ctx, key, key_len);
		else
			md5_update(&ctx, alt_digest, 16);

		/* Add salt for numbers not divisible by 3.  */
		if(cnt % 3 != 0)
			md5_update(&ctx, salt, salt_len);

		/* Add key for numbers not divisible by 7.  */
		if(cnt % 7 != 0)
			md5_update(&ctx, key, key_len);

		/* Add key or last result.  */
		if((cnt & 1) != 0)
			md5_update(&ctx, alt_digest, 16);
		else
			md5_update(&ctx, key, key_len);

		/* Create intermediate result.  */
		md5_final(alt_digest, &ctx);
	}

	for(cnt = 0; cnt < 16; ++cnt)
		digest[cnt] = alt_digest[cnt];
}

/* Kernel for the dictionary breaking mode: each thread analyzes a block of the dictionary,
   saving all the hashes it computes in hashes. */
#if DEBUG
__global__ void compute_dictionary(int size, u8 *hashes, grid_size gs, int len, u8 *dict, u8 *salt, u8 *target_hash, u8 *result_key) {
#else
__global__ void compute_dictionary(int size, grid_size gs, int len, u8 *dict, u8 *salt, u8 *target_hash, u8 *result_key) {
#endif
	int id, index, word_index, equal;
	u8 *string, *c;
	u8 digest[16];

	id = blockIdx.x * blockDim.x + threadIdx.x;
	if(id >= (gs.threads*gs.blocks))
		return;
	/* Identify current word index. */
	word_index = 0;
	c = dict + id*size;
	if(id != 0 && *c != '\0')
		++word_index;
	while(c != dict) {
		if(*c == '\0')
			++word_index;
		--c;
	}
	/* Identify starting string. */
	if(id == 0) {
		string = dict;
	} else {
		string = dict + id*size;
		while(*string != '\0') {
			++string;
		}
		++string;
	}
	/* Main body: analyze all the words in the dictionary block. */
	while(string <= dict + (id + 1)*size && *result_key == NULL) {
		#if DEBUG
		for(index = 0; index < 16; ++index)
			digest[index] = 0x00;
		#endif
		compute_md5(string, salt, digest);
		#if DEBUG
		for(index = 0; index < 16; ++index)
			hashes[word_index*16 + index] = digest[index];
		#endif
		/* Compare hashes. */
		equal = 1;
		for(index = 0; index < 16; ++index) {
			if(digest[index] != target_hash[index]) {
				equal = 0;
				break;
			}
		}
		if(equal) {
			custrcpy(result_key, string, custrlen(string));
		}
		while(*string != '\0')
			++string;
		++string;
		++word_index;
	}
	/* The last thread finishes the dictionary. */
	if(id == (gs.threads*gs.blocks - 1) && *result_key == '\0') {
		while(string < dict + len && *result_key == '\0') {
			#if DEBUG
			for(index = 0; index < 16; ++index)
				digest[index] = 0x00;
			#endif
			compute_md5(string, salt, digest);
			#if DEBUG
			for(index = 0; index < 16; ++index)
				hashes[word_index*16 + index] = digest[index];
			#endif
			/* Compare hashes. */
			equal = 1;
			for(index = 0; index < 16; ++index) {
				if(digest[index] != target_hash[index]) {
					equal = 0;
					break;
				}
			}
			if(equal) {
				custrcpy(result_key, string, custrlen(string));
			}
			while(*string != '\0')
				++string;
			++string;
			++word_index;
		}
	}
}

/* Kernel for the brute force mode: each thread generates and analyzes a group of passwords,
   saving all the hashes it computes in hashes. */
#if DEBUG
__global__ void compute_brute_force(u64 size, u8 *hashes, grid_size gs, bruteforce_size bfs, u8 *strings, u8 *salt, u8 *target_hash, u8 *result_key) {
#else
__global__ void compute_brute_force(u64 size, grid_size gs, bruteforce_size bfs, u8 *salt, u8 *target_hash, u8 *result_key) {
#endif
	u8 alphabet[] = ALPHABET;
	u8 alphaindex[MAX_PASSW_LENGTH];
	u8 string[MAX_PASSW_LENGTH];
	int id, exp, remainder, i, index, equal;
	u64 int_string, base, pow;
	u8 digest[16];
	u8 spill;

	id = blockIdx.x * blockDim.x + threadIdx.x;
	if(id >= gs.blocks*gs.threads)
		return;
	int_string = id*size;

	/* Compute first string from int_string. */
	base = int_string;
	for(i = 0; i < MAX_PASSW_LENGTH; ++i) {
		string[i] = '\0';
	}
	for(exp = bfs.len - 1; exp >= 0; --exp) {
		pow = 1;
		for(i = 0; i < exp; ++i) {
			pow *= ALPHALEN;
		}
		remainder = base/pow;
		string[bfs.len - 1 - exp] = alphabet[remainder];
		alphaindex[bfs.len - 1 - exp] = remainder;
		base -= remainder * pow;
	}
	string[bfs.len - 1 - exp] = '\0';

	while(int_string < (id + 1)*size && *result_key == '\0') {
		#if DEBUG
		/* Copy generated string into string memory. */
		custrcpy(strings + int_string*MAX_PASSW_LENGTH, string, custrlen(string));
		/* Reset digest. */
		for(index = 0; index < 16; ++index)
			digest[index] = 0x0;
		#endif
		/* Compute md5 and store it into hashes memory. */
		compute_md5(string, salt, digest);
		#if DEBUG
		for(index = 0; index < 16; ++index)
		 	hashes[int_string*16 + index] = digest[index];
		#endif
		/* Compare computed hash with target hash. */
		equal = 1;
		for(index = 0; index < 16; ++index) {
			if(digest[index] != target_hash[index]) {
				equal = 0;
				break;
			}
		}
		/* Eventually copy string corresponding to correct hash into result key. */
		if(equal) {
			custrcpy(result_key, string, custrlen(string));
		}
		++int_string;
		/* Generate next password incrementally from last string. */
		spill = 1;
		index = bfs.len - 1;
		while(spill && index >= 0) {
			spill = 0;
			++alphaindex[index];
			if(alphaindex[index] >= ALPHALEN) {
				spill = 1;
				alphaindex[index] = 0;
				--index;
			}
		}
		/* Generate actual string from alphaindex. */
		for(index = 0; index < bfs.len; ++index)
			string[index] = alphabet[alphaindex[index]];
	}
	/* The last thread finishes to analyze the key space. */
	if(id == ((gs.blocks*gs.threads) - 1) && *result_key == '\0') {
		while(int_string < bfs.alphabet_size && *result_key == '\0') {
			/* Generate current password from int_string. */
			base = int_string;
			for(i = 0; i < MAX_PASSW_LENGTH; ++i) {
				string[i] = '\0';
			}
			for(exp = bfs.len - 1; exp >= 0; --exp) {
				pow = 1;
				for(i = 0; i < exp; ++i) {
					pow *= ALPHALEN;
				}
				remainder = base/pow;
				string[bfs.len - 1 - exp] = alphabet[remainder];
				base -= remainder * pow;
			}
			string[bfs.len - 1 - exp] = '\0';
			#if DEBUG
			/* Copy generated string into string memory. */
			custrcpy(strings + int_string*MAX_PASSW_LENGTH, string, custrlen(string));
			/* Reset digest. */
			for(index = 0; index < 16; ++index)
				digest[index] = 0x00;
			#endif
			/* Compute md5 and store it into hashes memory. */
			compute_md5(string, salt, digest);
			#if DEBUG
			for(index = 0; index < 16; ++index)
				hashes[int_string*16 + index] = digest[index];
			#endif
			/* Compare computed hash with target hash. */
			equal = 1;
			for(index = 0; index < 16; ++index) {
				if(digest[index] != target_hash[index]) {
					equal = 0;
					break;
				}
			}
			/* Eventually copy string corresponding to correct hash into result key. */
			if(equal) {
				custrcpy(result_key, string, custrlen(string));
			}
			++int_string;
		}
	}
}

/* Function that prepares the memory on the GPU and calls the dictionary kernel function. If DEBUG == 1 it then prints the computed hashes. */
void break_dictionary(grid_size gs, dict_size ds, u8 *dict, u8 *salt, u8 *target_hash, u8 *result_key) {
	#if DEBUG
	u8 *hashes; // It is used to store the generated hashes on the GPU.
	u8 *host_hashes; // It is used to store and print the hashes.
	int i; // Index used to print the computed hashes.
	#endif
	int index; // Index used in loops.
	int size; // Bytes of dictionary that each thread has to analyze.
	u8 *host_result_key; // It is used to store and print the eventual result key.
	md5_ctx host_init_ctx; // Initial MD5 context.
	
	/* Create MD5 initial context and copy it to constant memory. */
	host_init_ctx.total[0] = 0;
	host_init_ctx.total[1] = 0;
	host_init_ctx.A = 0x67452301;
	host_init_ctx.B = 0xefcdab89;
	host_init_ctx.C = 0x98badcfe;
	host_init_ctx.D = 0x10325476;
	host_init_ctx.buflen = 0;
	for(index = 0; index < 128; ++index)
		host_init_ctx.buffer[index] = 0;
	printf("Memcpy %d\n", cudaMemcpyToSymbol(init_ctx, &host_init_ctx, sizeof(md5_ctx), 0, cudaMemcpyHostToDevice));

	/* Compute per thread job. */
	size = ds.len/(gs.blocks*gs.threads);
	#if DEBUG
	/* Allocate memory for the computed hashes. */
	cudaMalloc((void**) &hashes, sizeof(u8)*16*ds.word_cnt);
	cudaMemset(hashes, 0, sizeof(u8)*16*ds.word_cnt);
	compute_dictionary<<<gs.blocks, gs.threads>>>(size, hashes, gs, ds.len, dict, salt, target_hash, result_key);
	#else
	compute_dictionary<<<gs.blocks, gs.threads>>>(size, gs, ds.len, dict, salt, target_hash, result_key);
	#endif

	/* Allocate memory for printing result key. */
	host_result_key = (u8*) malloc(sizeof(u8)*MAX_PASSW_LENGTH);
	cudaMemcpy(host_result_key, result_key, sizeof(u8)*MAX_PASSW_LENGTH, cudaMemcpyDeviceToHost);
	#if DEBUG
	/* Allocate memory for printing computed hashes. */
	host_hashes = (u8*) malloc(sizeof(u8)*16*ds.word_cnt);
	cudaMemcpy(host_hashes, hashes, sizeof(u8)*16*ds.word_cnt, cudaMemcpyDeviceToHost);
	/* Print computed hashes. */
	fprintf(stderr, "Printing hashes:\n");
	for(index = 0; index < ds.word_cnt; ++index) {
		for(i = 0; i < 16; ++i) {
			fprintf(stderr, "%02x", host_hashes[index*16 + i]);
		}
		fprintf(stderr, "\n");
	}
	#endif
	printf("Right key: %s\n", host_result_key);

	#if DEBUG
	cudaFree(hashes);
	free(host_hashes);
	#endif
	free(host_result_key);
}

/* Function that prepares the memory on the GPU and calls the brute force kernel function. If DEBUG == 1 it then prints the generated passwords and the corresponding hashes. */
void brute_force(int len, grid_size gs, u8 *salt, u8 *target_hash, u8 *result_key) {
	#if DEBUG
	u8 *hashes; // It is used to store the generated hashes on the GPU.
	u8 *host_hashes; // It is used to store and print the hashes.
	u8 *strings; // It is used to store the generated passwords on the GPU.
	u8 *host_strings; // It is used to store and print the passwords.
	#endif
	u64 index; // Index used in loops.
	u64 size; // Number of passwords that each thread has to analyze.
	u8 i; // General purpose index.
	u8 *host_result_key; // It is used to store and print the eventual result key.
	bruteforce_size bfs; // Holds alphabet size and password length.
	md5_ctx host_init_ctx; // Initial MD5 context.
	
	/* Create MD5 initial context and copy it to constant memory. */
	host_init_ctx.total[0] = 0;
	host_init_ctx.total[1] = 0;
	host_init_ctx.A = 0x67452301;
	host_init_ctx.B = 0xefcdab89;
	host_init_ctx.C = 0x98badcfe;
	host_init_ctx.D = 0x10325476;
	host_init_ctx.buflen = 0;
	for(index = 0; index < 128; ++index)
		host_init_ctx.buffer[index] = 0;
	printf("Memcpy %d\n", cudaMemcpyToSymbol(init_ctx, &host_init_ctx, sizeof(md5_ctx), 0, cudaMemcpyHostToDevice));

	/* Set password length. */
	bfs.len = len;

	/* Compute alphabet size. */
	bfs.alphabet_size = 1;
	for(i = 0; i < len; ++i) {
		bfs.alphabet_size *= ALPHALEN;
	}
	/* Compute per thread job. */
	size = bfs.alphabet_size/(gs.threads*gs.blocks);
	#if DEBUG
	/* Allocate memory for computed hashes. */
	cudaMalloc((void**) &hashes, sizeof(u8)*16*bfs.alphabet_size);
	cudaMemset(hashes, 0, sizeof(u8)*16*bfs.alphabet_size);
	/* Allocate memory for generated passwords. */
	cudaMalloc((void**) &strings, sizeof(u8)*MAX_PASSW_LENGTH*bfs.alphabet_size);
	cudaMemset(strings, '\0', sizeof(u8)*MAX_PASSW_LENGTH*bfs.alphabet_size);
	fprintf(stderr, "size %lu, alphabet_size %lu\n", size, bfs.alphabet_size);
	/* Invoke brute force kernel. */
	compute_brute_force<<<gs.blocks, gs.threads>>>(size, hashes, gs, bfs, strings, salt, target_hash, result_key);
	#else
	compute_brute_force<<<gs.blocks, gs.threads>>>(size, gs, bfs, salt, target_hash, result_key);
	#endif
	/* Allocate memory for result key and copy it from GPU. */
	host_result_key = (u8*) malloc(sizeof(u8)*MAX_PASSW_LENGTH);
	cudaMemcpy(host_result_key, result_key, sizeof(u8)*MAX_PASSW_LENGTH, cudaMemcpyDeviceToHost);

	#if DEBUG
	/* Allocate memory for printing hashes. */
	host_hashes = (u8*) malloc(sizeof(u8)*16*bfs.alphabet_size);
	cudaMemcpy(host_hashes, hashes, sizeof(u8)*16*bfs.alphabet_size, cudaMemcpyDeviceToHost);
	/* Allocate memory for printing passwords. */
	host_strings = (u8*) malloc(sizeof(u8)*MAX_PASSW_LENGTH*bfs.alphabet_size);
	cudaMemcpy(host_strings, strings, sizeof(u8)*MAX_PASSW_LENGTH*bfs.alphabet_size, cudaMemcpyDeviceToHost);
	/* Print computed hashes and passwords. */
	fprintf(stderr, "Printing hashes:\n");
	for(index = 0; index < bfs.alphabet_size; ++index) {
		fprintf(stderr, "%s :", &host_strings[index*MAX_PASSW_LENGTH]);
		for(i = 0; i < 16; ++i) {
			fprintf(stderr, "%02x", host_hashes[index*16 + i]);
		}
		fprintf(stderr, "\n");
	}
	#endif
	fprintf(stderr, "Right key: %s\n", host_result_key);

	#if DEBUG
	cudaFree(hashes);
	cudaFree(strings);
	free(host_hashes);
	free(host_strings);
	#endif
	free(host_result_key);
}
